// phalanx config.hpp.cu header file

// Copyright (c) 2011 - 2011 Kohei Takahashi (Flast).
// Distributed under the MIT license. for more detail see COPYING.

#ifndef IG_PHALANX_DETAIL_CONFIG_HPP_CU_ONCE_
#define IG_PHALANX_DETAIL_CONFIG_HPP_CU_ONCE_

#include <boost/preprocessor/control/if.hpp>
#include <boost/preprocessor/logical/and.hpp>

#include <phalanx/detail/version.hpp.cu>

#define LO_PHALANX_CONFIG_ATTRIBUTE_host        __host__
#define LO_PHALANX_CONFIG_ATTRIBUTE_global      __global__
#define LO_PHALANX_CONFIG_ATTRIBUTE_device      __device__

#if PHALANX_CUDA_WORKAROUND( <= 3010 )
#define LO_PHALANX_CONFIG_ATTRIBUTE_static_device __device__
#else
#define LO_PHALANX_CONFIG_ATTRIBUTE_static_device static __device__
#endif

#define LO_PHALANX_CONFIG_ATTRIBUTE_shared      __shared__
#define LO_PHALANX_CONFIG_ATTRIBUTE_constant    __constant__
#define LO_PHALANX_CONFIG_ATTRIBUTE_restrict    __restrict__

#define LO_PHALANX_CONFIG_ATTRIBUTE_inline      inline
#define LO_PHALANX_CONFIG_ATTRIBUTE_forceinline __forceinline__
#define LO_PHALANX_CONFIG_ATTRIBUTE_noinline    __noinline__

#define PHALANX_CONFIG_ATTRIBUTE( attribute ) \
  LO_PHALANX_CONFIG_ATTRIBUTE_ ## attribute

#define PHALANX_CONFIG_DEVICE_FUNCTION \
  PHALANX_CONFIG_ATTRIBUTE( inline ) \
  PHALANX_CONFIG_ATTRIBUTE( device )

#define PHALANX_CONFIG_STATIC_DEVICE_FUNCTION \
  PHALANX_CONFIG_ATTRIBUTE( inline ) \
  PHALANX_CONFIG_ATTRIBUTE( static_device )

#define PHALANX_CONFIG_ANY_FUNCTION \
  PHALANX_CONFIG_DEVICE_FUNCTION \
  PHALANX_CONFIG_ATTRIBUTE( host )

#define PHALANX_CONFIG_STATIC_ANY_FUNCTION \
  PHALANX_CONFIG_STATIC_DEVICE_FUNCTION \
  PHALANX_CONFIG_ATTRIBUTE( host )

#if !defined( PHALANX_CONFIG_DEFAULT_NOT_TO_ANY_FUNCTION )
#define LO_PHALANX_CONFIG_DEFAULT_IS_ANY 1
#else
#define LO_PHALANX_CONFIG_DEFAULT_IS_ANY 0
#endif

#define LO_PHALANX_CONFIG_BOOL_true  1
#define LO_PHALANX_CONFIG_BOOL_false 0
#define LO_PHALANX_CONFIG_BOOL( pred ) LO_PHALANX_CONFIG_BOOL_ ## pred

#define PHALANX_CONFIG_DEFAULT_FUNCTION( pred ) \
  BOOST_PP_IF( \
    BOOST_PP_AND( \
      LO_PHALANX_CONFIG_BOOL( pred ), \
      LO_PHALANX_CONFIG_DEFAULT_IS_ANY ), \
    PHALANX_CONFIG_ANY_FUNCTION, \
    PHALANX_CONFIG_DEVICE_FUNCTION )

#define PHALANX_CONFIG_DEFAULT_STATIC_FUNCTION( pred ) \
  BOOST_PP_IF( \
    BOOST_PP_AND( \
      LO_PHALANX_CONFIG_BOOL( pred ), \
      LO_PHALANX_CONFIG_DEFAULT_IS_ANY ), \
    PHALANX_CONFIG_STATIC_ANY_FUNCTION, \
    PHALANX_CONFIG_STATIC_DEVICE_FUNCTION )

#endif // IG_PHALANX_DETAIL_CONFIG_HPP_CU_ONCE_

